#ifndef REDUCE_H
#define REDUCE_H
#include "reductionOperators.h"

///
//! take one step for Summation
///
template<int stepSize>
__device__ void sumStep(float* temps) {
	if (threadIdx.x < stepSize) {
		temps[threadIdx.x] += temps[threadIdx.x + stepSize];
	}
	if (stepSize >= 32) { //from hardware side, 1 warp has 32 threads, 
		__syncthreads();  //when step size smaller than 32, do not need sync
	}
}

///
//! take 9 steps for Summation
//! use template to generate 9 sumStep() function
///
__device__ static void sumReduce(float* temps) {
	if (256 < BLOCKSIZE) sumStep<256>(temps);
	if (128 < BLOCKSIZE) sumStep<128>(temps);
	if ( 64 < BLOCKSIZE) sumStep< 64>(temps);
	if ( 32 < BLOCKSIZE) sumStep< 32>(temps);
	if ( 16 < BLOCKSIZE) sumStep< 16>(temps);
	if (  8 < BLOCKSIZE) sumStep<  8>(temps);
	if (  4 < BLOCKSIZE) sumStep<  4>(temps);
	if (  2 < BLOCKSIZE) sumStep<  2>(temps);
	if (  1 < BLOCKSIZE) sumStep<  1>(temps);
}


///
//! one step for finding Maximum
///
template<int stepSize>
__device__ void maxStep(float* temps) {
	if (threadIdx.x < stepSize) {
		maxOperator(temps[threadIdx.x], temps[threadIdx.x + stepSize], temps + threadIdx.x);
	}
	if (stepSize >= 32) { //from hardware side, 1 warp has 32 threads, 
		__syncthreads();  //when step size smaller than 32, do not need sync
	}
}

///
//! take 9 steps for finding Maximum
//! use template to generate 9 maxStep() function, when 9 step finish, the maximum is stored in temps[0]
///
__device__ static void maxReduce(float* temps) {
	if (256 < BLOCKSIZE) maxStep<256>(temps);
	if (128 < BLOCKSIZE) maxStep<128>(temps);
	if ( 64 < BLOCKSIZE) maxStep< 64>(temps);
	if ( 32 < BLOCKSIZE) maxStep< 32>(temps);
	if ( 16 < BLOCKSIZE) maxStep< 16>(temps);
	if (  8 < BLOCKSIZE) maxStep<  8>(temps);
	if (  4 < BLOCKSIZE) maxStep<  4>(temps);
	if (  2 < BLOCKSIZE) maxStep<  2>(temps);
	if (  1 < BLOCKSIZE) maxStep<  1>(temps);
}


///
//! take one step for finding Minimum index-value pair
///
template<int stepSize>
__device__ static void argminStep(float* values, int* indices) {
	if (threadIdx.x < stepSize) {
		int compOffset = threadIdx.x + stepSize;
		argMin(indices[threadIdx.x], values[threadIdx.x], indices[compOffset], values[compOffset], indices + threadIdx.x, values + threadIdx.x);
	}
	if (stepSize >= 32) { //from hardware side, 1 warp has 32 threads, 
		__syncthreads();  //when step size smaller than 32, do not need sync
	}
}

///
//! take 9 steps for finding Minimum index-value pair
//! use template to generate 9 argminStep() function, when 9 step finish, the minimum pair is stored in indices[0] & values[0]
///
__device__ static void argminReduce(float* values, int* indices) {
	if (256 < BLOCKSIZE) argminStep<256>(values, indices);
	if (128 < BLOCKSIZE) argminStep<128>(values, indices);
	if ( 64 < BLOCKSIZE) argminStep< 64>(values, indices);
	if ( 32 < BLOCKSIZE) argminStep< 32>(values, indices);
	if ( 16 < BLOCKSIZE) argminStep< 16>(values, indices);
	if (  8 < BLOCKSIZE) argminStep<  8>(values, indices);
	if (  4 < BLOCKSIZE) argminStep<  4>(values, indices);
	if (  2 < BLOCKSIZE) argminStep<  2>(values, indices);
	if (  1 < BLOCKSIZE) argminStep<  1>(values, indices);
}


///
/// take one step for finding Maximum index-value pair
///
template<int stepSize>
__device__ void argmaxStep(float* values, int* indices) {
	if (threadIdx.x < stepSize) {
		int compOffset = threadIdx.x + stepSize;
		argMax(indices[threadIdx.x], values[threadIdx.x], indices[compOffset], values[compOffset], indices + threadIdx.x, values + threadIdx.x);
	}
	if (stepSize >= 32) { //from hardware side, 1 warp has 32 threads, 
		__syncthreads();  //when step size smaller than 32, do not need sync
	}
}

///
/// take 9 step for finding Maximum
/// use template to generate 9 argmaxStep() function, when 9 steps finish, the maximum pair is stored in indices[0] & values[0]
///
__device__ static void argmaxReduce(float* values, int* indices) {
	if (256 < BLOCKSIZE) argmaxStep<256>(values, indices);
	if (128 < BLOCKSIZE) argmaxStep<128>(values, indices);
	if ( 64 < BLOCKSIZE) argmaxStep< 64>(values, indices);
	if ( 32 < BLOCKSIZE) argmaxStep< 32>(values, indices);
	if ( 16 < BLOCKSIZE) argmaxStep< 16>(values, indices);
	if (  8 < BLOCKSIZE) argmaxStep<  8>(values, indices);
	if (  4 < BLOCKSIZE) argmaxStep<  4>(values, indices);
	if (  2 < BLOCKSIZE) argmaxStep<  2>(values, indices);
	if (  1 < BLOCKSIZE) argmaxStep<  1>(values, indices);
}

#endif
